#include<iostream>
#include<cuda_runtime.h>
#include "../MyTimer.h"


__device__ float devData; //设备的全局变量。


int main()
{
    float val = 1;
    cudaMemcpyToSymbol(devData, &val, sizeof(float));
    cudaMemcpyFromSymbol(&val, devData,sizeof(float));

    //或者
    float *devPtr;
    cudaGetSymbolAddress((void**)&devPtr, devData);
    //对于主机端的设备变量是不能直接&devData的。
    cudaMemcpy(devPtr, val, sizeof(float), cudaMemcpyHostToDevice);


    //cudaHostAlloc:创建一个到固定内存的映射。用直接new的话系统页可交换，可能导致数据传输慢	
    /*
    flag:
    ·cudaHostAllocDefalt 内存分配在主机上，页锁定（不能被换出内存页）只是想提高主机到设备传输的速度，用这个就够了。
    ·cudaHostAllocPortable 分配的内存在多GPU 系统中是共享的；多 GPU 编程时用它，内存共享更方便。
    ·cudaHostAllocWriteCombined 写合并（Write-Combining）意味着 CPU 向内存写入是合并写、但不能高效读取；适用于“主机写 → GPU 读”的单向数据流（只写不读）
    ·cudaHostAllocMapped 零拷贝内存的最明显的标志是cudaHostAllocMapped，分配的主机内存会自动映射到设备端；✅ 在小数据 + 低延迟传输场景下很好用（如：GPU 直接读取少量主机参数）
    cudaHostGetDevicePointer获取映射到固定内存的设备指针


    float *h_ptr;
    float *d_ptr;

    cudaSetDeviceFlags(cudaDeviceMapHost);
    cudaHostAlloc((void**)&h_ptr, size, cudaHostAllocMapped);
    cudaGetDevicePointer((void**)&d_ptr, h_ptr, 0);

    kernel<<<1, 1>>>(d_ptr);  // GPU 直接访问 host 分配的内存


    具体例子
    float *h_data = nullptr;  // 主机数据指针（页锁定）
    float *d_data = nullptr;  // 设备数据指针
    // 分配页锁定主机内存（Pinned memory）
    cudaHostAlloc((void**)&h_data, size, cudaHostAllocDefault);
    // 初始化主机数据
    for (int i = 0; i < N; ++i) {
        h_data[i] = static_cast<float>(i);
    }
    // 分配设备内存
    cudaMalloc((void**)&d_data, size);
    // 复制主机数据到设备
    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
    // 启动 kernel
    addOne<<<1, N>>>(d_data, N);
    // 复制结果回主机
    cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);


    
    */


}











